#include "opencl_source_map.hpp" 
namespace MNN { 
const char* winogradTransformSource2_3_1 = 
"#ifdef MNN_SUPPORT_FP16\n"
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
"#endif\n"
"__constant sampler_t SAMPLER=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n"
"__kernel void winogradTransformSource(__read_only image2d_t uInput,// 0\n"
" __write_only image2d_t uOutput,__private const int unitWidth,\n"
" __private const int unitHeight,// 3\n"
" __private const int padX,__private const int padY,\n"
" __private const int srcWidth,// 6\n"
" __private const int srcHeight,__private const int srcChannelC4,\n"
" __private const int batchOffset) {\n"
" int2 pos=(int2)(get_global_id(0),get_global_id(1));\n"
" if (pos.x<unitWidth*unitHeight && pos.y<srcChannelC4) {\n"
" int unitWidth_idx=pos.x % unitWidth;\n"
" int unitHeight_idx=pos.x/unitWidth;\n"
" int dstX=mad24(pos.y,unitWidth,unitWidth_idx);\n"
" {\n"
" int sxStart=unitWidth_idx*2-padX;\n"
" int syStart=unitHeight_idx*2-padY;\n"
" FLOAT4 S00;\n"
" FLOAT4 S10;\n"
" FLOAT4 S20;\n"
" FLOAT4 S30;\n"
" FLOAT4 S01;\n"
" FLOAT4 S11;\n"
" FLOAT4 S21;\n"
" FLOAT4 S31;\n"
" FLOAT4 S02;\n"
" FLOAT4 S12;\n"
" FLOAT4 S22;\n"
" FLOAT4 S32;\n"
" FLOAT4 S03;\n"
" FLOAT4 S13;\n"
" FLOAT4 S23;\n"
" FLOAT4 S33;\n"
" {\n"
" int sx=0+sxStart;\n"
" int sy=0+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S00=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=1+sxStart;\n"
" int sy=0+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S10=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=2+sxStart;\n"
" int sy=0+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S20=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=3+sxStart;\n"
" int sy=0+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S30=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=0+sxStart;\n"
" int sy=1+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S01=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=1+sxStart;\n"
" int sy=1+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S11=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=2+sxStart;\n"
" int sy=1+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S21=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=3+sxStart;\n"
" int sy=1+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S31=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=0+sxStart;\n"
" int sy=2+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S02=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=1+sxStart;\n"
" int sy=2+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S12=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=2+sxStart;\n"
" int sy=2+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S22=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=3+sxStart;\n"
" int sy=2+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S32=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=0+sxStart;\n"
" int sy=3+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S03=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=1+sxStart;\n"
" int sy=3+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S13=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=2+sxStart;\n"
" int sy=3+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S23=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" {\n"
" int sx=3+sxStart;\n"
" int sy=3+syStart;\n"
" int imageSx=select(sx+pos.y*srcWidth,-1,sx<0 || sx >= srcWidth);\n"
" int imageSy=select(batchOffset*srcHeight+sy,-1,sy<0 || sy >= srcHeight);\n"
" S33=RI_F(uInput,SAMPLER,(int2)(imageSx,imageSy));\n"
" }\n"
" FLOAT4 m00=+S00-S02;\n"
" FLOAT4 m10=+S10-S12;\n"
" FLOAT4 m20=+S20-S22;\n"
" FLOAT4 m30=+S30-S32;\n"
" FLOAT4 m01=+(FLOAT)0.5f*S01+(FLOAT)0.5f*S02;\n"
" FLOAT4 m11=+(FLOAT)0.5f*S11+(FLOAT)0.5f*S12;\n"
" FLOAT4 m21=+(FLOAT)0.5f*S21+(FLOAT)0.5f*S22;\n"
" FLOAT4 m31=+(FLOAT)0.5f*S31+(FLOAT)0.5f*S32;\n"
" FLOAT4 m02=-(FLOAT)0.5f*S01+(FLOAT)0.5f*S02;\n"
" FLOAT4 m12=-(FLOAT)0.5f*S11+(FLOAT)0.5f*S12;\n"
" FLOAT4 m22=-(FLOAT)0.5f*S21+(FLOAT)0.5f*S22;\n"
" FLOAT4 m32=-(FLOAT)0.5f*S31+(FLOAT)0.5f*S32;\n"
" FLOAT4 m03=-S01+S03;\n"
" FLOAT4 m13=-S11+S13;\n"
" FLOAT4 m23=-S21+S23;\n"
" FLOAT4 m33=-S31+S33;\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*0),+m00-m20);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*1),+(FLOAT)0.5f*m10+(FLOAT)0.5f*m20);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*2),-(FLOAT)0.5f*m10+(FLOAT)0.5f*m20);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*3),-m10+m30);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*4),+m01-m21);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*5),+(FLOAT)0.5f*m11+(FLOAT)0.5f*m21);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*6),-(FLOAT)0.5f*m11+(FLOAT)0.5f*m21);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*7),-m11+m31);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*8),+m02-m22);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*9),+(FLOAT)0.5f*m12+(FLOAT)0.5f*m22);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*10),-(FLOAT)0.5f*m12+(FLOAT)0.5f*m22);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*11),-m12+m32);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*12),+m03-m23);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*13),+(FLOAT)0.5f*m13+(FLOAT)0.5f*m23);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*14),-(FLOAT)0.5f*m13+(FLOAT)0.5f*m23);\n"
" WI_F(uOutput,(int2)(dstX,unitHeight_idx+unitHeight*15),-m13+m33);\n"
" }\n"
" }\n"
"}\n"
;
}
